#ifndef	gating_included
#define	gating_included

__global__	void	_gatingf(uint32_t	R,	half	*u,	half	*out) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3,	r=id%R,	c=id/R;
    *(half8*)(out+id)=_hmul8(*(half8*)(u+c*2*R+r),*(half8*)(u+c*2*R+R+r));
}

__global__	void	_gatingb(uint32_t	R,	half	*u,	half	*gin,	half	*U) {
    uint32_t	id=(blockIdx.x*blockDim.x+threadIdx.x)<<3,	r=id%R,	c=id/R;
    *(half8*)(U+c*2*R+r)=_hmul8(*(half8*)(u+c*2*R+R+r),*(half8*)(gin+id));
    *(half8*)(U+c*2*R+R+r)=_hmul8(*(half8*)(u+c*2*R+r),*(half8*)(gin+id));
}

template<uint32_t	R,	uint32_t	C>
struct	gating {
    void	fw(tensor<half,2*R*C>	&inp/*const*/,tensor<half,R*C>	&out/*changed*/) {
        uint32_t	w=warp_size(8,R*C);
        _gatingf<<<R*C/8/w,w>>>(R,inp.p,out.p);
    }

    void	bk(tensor<half,2*R*C>	&inp/*const*/,tensor<half,R*C>	&gin/*const*/,tensor<half,2*R*C>	&gra/*changed*/) {
        uint32_t	w=warp_size(8,R*C);
        _gatingb<<<R*C/8/w,w>>>(R,inp.p,gin.p,gra.p);
    }
};

#endif

